#include "GpuOperation.h"

/************************************************************************/
/* texture fetcher                                                               */
/************************************************************************/
texture<float, 1, cudaReadModeElementType> myTex; ///cached fetcher

/************************************************************************/
/* kernel                                                               */
/************************************************************************/
__global__ void incrementArrayOnDevice(float* result, int N)
{
	int idx = blockIdx.x * blockDim.x + threadIdx.x;
	if (idx < N) 
		result[idx] = tex1Dfetch(myTex, idx) + 1.0f;
};

void bindTexture( CH::GlobalMemory<float> & a_d)
{
	a_d.bind_texture( myTex ) ;
}

void launchKernel( CH::GlobalMemory<float> & r_d)
{
	int blockSize = 4;
    int nBlocks = r_d.size() / blockSize + (r_d.size() % blockSize == 0?0:1);
	incrementArrayOnDevice <<< nBlocks, blockSize >>> (r_d.get_memory_ptr(), r_d.size());
}

